feat: rebuild the extension as a production-grade CUDA driver (v0.2.0) - #3
Merged
Conversation
l4nos
force-pushed
the
feat/v0.2.0-driver-rebuild
branch
3 times, most recently
from
September 5, 2026 19:42
3c8cfc6 to
81abe41
Compare
## Summary This is a ground-up rebuild of the PHP CUDA extension into the driver layer that makes a "PHPTorch" technically possible: device-resident tensors with a real memory model, the NVIDIA library stack bound with correct semantics, streams/events/CUDA graphs, and NVRTC runtime kernel compilation, on a build system that actually compiles on modern toolchains. It also fixes every root cause of issue #1 (Ubuntu 24.04 + CUDA 12 + cuDNN 9 build failure). 50 files changed. Fixes #1 --- ## What was broken (found in a full code review) The tree did not compile at all, on any platform: - `cuda.c` referenced module globals (`allow_async_operations`, `default_device`) that did not exist in the globals struct: hard compile error - 7 functions in the Zend function table had no implementations: link error - `cuda_matrix_multiply_kernel_wrapper()` was declared extern, defined nowhere - ~50 `PHP_FUNCTION` declarations in `php_cuda.h` had no implementations at all, while the test suite called several of them - `compile.sh` deleted `config.m4` via a `config.*` glob on every run, so the second build destroyed the project's build configuration - `config.m4` had no working rule to compile `.cu` files (`Makefile.frag` was never included and referenced undefined variables) - `memory_pool.cuh` used `std::atomic` without `#include <atomic>` Issue #1 root causes specifically: - `-arch=sm_30`: Kepler support was removed in CUDA 12; nvcc 12.x rejects it - `cudnnGetConvolutionForwardAlgorithm`: removed in cuDNN 9 - NVTX on CUDA 12 is header-only (`nvtx3/`); the build required `-lnvToolsExt` - Reporter's environment had a mixed toolchain (nvcc 12.0 vs toolkit 12.9) and a driver (535, max runtime 12.2) older than the toolkit runtime Runtime correctness bugs: - `cuda_free()` double-freed (explicit `cudaFree` + resource destructor) - `cuda_matrix_multiply()` had wrong dimension math (only worked for square matrices), no inner-dimension validation, and read an uninitialized variable - Row-major PHP data was fed to column-major cuBLAS as-is (wrong results) - `neural_net.cu` had an out-of-bounds read and null derefs - `cudaDeviceReset()` ran unconditionally at module shutdown, which is hostile under php-fpm and persistent workers - MINIT required a GPU, making the extension unloadable on GPU-less hosts --- ## What this commit delivers ### Build system (fixes #1) - `config.m4` rewritten: `.cu` compilation via `PHP_ADD_MAKEFILE_FRAGMENT` with working NVCC rules; GPU architectures auto-detected from `nvidia-smi` with toolkit-gated fallback lists (sm_70 through sm_120); `--with-cuda-arch=` override; cuDNN 7/8/9 detection with `CUDNN_MAJOR` gating; NVTX3 (header-only) vs legacy nvToolsExt detection; NVRTC + driver API (`-lnvrtc -lcuda`) with stubs path handling; OpenMP optional - `Makefile.frag` rewritten: NVCC pattern rules, module-target prerequisites - `compile.sh` rewritten: never touches `config.m4`; refuses mixed toolchains (nvcc vs `version.json`); warns when the driver-supported runtime is older than the toolkit; `--uninstall`; `--test`; PHP version check (8.1+) - `Dockerfile`: Ubuntu 24.04 + CUDA 12.6 + cuDNN 9 reference environment - GitHub Actions: compile-only matrix PHP {8.1-8.4} x CUDA {11.8, 12.6} plus a self-hosted GPU runner job for the `.phpt` suite; header self-containment check ### Core runtime - All correctness bugs above fixed - Complete core API: device management, malloc/free/memset, all memcpy directions (binary-safe strings), pinned + unified memory, memory info, bandwidth measurement, profiler start/stop - Lazy device initialization: the extension now loads on GPU-less hosts - Error model: `E_WARNING` + `false` by default; `cuda.error_mode=exception` throws `CudaException` - ini settings: `cuda.default_device`, `cuda.error_mode`, `cuda.enable_cpu_fallback`, `cuda.enable_memory_pool` - Memory pool: shared-refcount lifetime (safe against request-shutdown destruction order), block reuse, stats ### CudaTensor substrate - `CudaTensor` PHP class: refcounted storage separate from views. `reshape`, `transpose` and `slice` are zero-copy views sharing one allocation (the ATen/c10 memory model) - Broadcasting elementwise ops (`add/sub/mul/div` + in-place variants), `matmul` via cuBLAS with the row-major transpose trick (numerically correct, verified against reference values in tests) - 6 dtypes: fp32, fp64, int32, fp16, bf16, int8 (templated kernel dispatch) - Activations (relu/sigmoid/tanh/exp/log/sqrt/gelu/neg), softmax, reductions (sum/mean/max/min) - NVRTC: `cuda_kernel_compile()` / `cuda_kernel_launch()`: custom CUDA kernels authored in PHP userland, compiled to PTX at runtime, launched with CudaTensor/int/float arguments on any stream ### Execution model - Streams: create/destroy/synchronize/query/wait-event - Events: start/stop timing pairs with elapsed-time queries - CUDA graphs: `cuda_graph_begin_capture` / `end_capture` (instantiate) / `launch` / `destroy`: capture once, replay with near-zero CPU overhead ### NVIDIA library stack (partial) - cuBLAS: handle resources, `cuda_cublas_matrix_multiply`, `cuda_cublas_gemm` (alpha/beta), `cuda_batch_gemm`, all row-major correct - cuDNN 8/9: `cuda_cudnn_convolution_forward` with `_v7` algorithm discovery (the legacy API was removed in cuDNN 9), workspace management ### Dead code removed `neural_net.*` (superseded by the CudaTensor substrate), `matrix_ops.*`, `tensor_ops.*`, `cuda_kernel.cu`, `cudnn_advanced.cuh`, `logger.cuh`: all contained unimplemented or incorrect code referenced by nothing. ### Tests, docs, packaging - Test suite rewritten against the real API (11 `.phpt` files), including numeric verification of matmul/GEMM/convolution/softmax/broadcasting, an NVRTC kernel test, and a CUDA graph capture/replay test; tests skip gracefully without a GPU or optional components - `docs/INSTALL-UBUNTU.md`: exact Ubuntu 24.04 package set (including the `nvidia-cuda-toolkit` conflict warning from #1) plus a troubleshooting table - README rewritten (documents only what exists) - `composer.json` with `php-ext` configure options for PIE --- ## Verification - `phpize` passes - All 6 C files pass `gcc -fsyntax-only` against PHP 8.3 headers - All 8 .cu files pass C++ syntax checking (4 real bugs found and fixed this way: missing `void**` cast, missing `<stdexcept>`, 1-arg `cudaEventRecord`, missing `cuda_fp16.h` include) - Function table, implementations and header declarations cross-checked: zero orphans; all 40 CudaTensor methods match their registration - All test PHP lints clean - GPU execution is the remaining gate: `./compile.sh --test` on a CUDA host or the provided Docker image (`docker run --gpus all`) ## Known limitations (follow-up work) - cuBLASLt, cuSOLVER, cuFFT, cuRAND, cuSPARSE, NCCL, NVML/CUPTI: planned - `CudaTensor` matmul is 2-D only (batched matmul planned) - DLPack interop and `fromBuffer` planned - Windows unsupported by design; macOS has no CUDA --- ## CI follow-up fixes (verified in the exact CI containers) The first CI run exposed that the module was not actually building on the matrix. All issues below were reproduced and fixed against nvidia/cuda:12.6.3-cudnn-devel-ubuntu24.04 and PHP 8.1/8.4 locally: - `config.m4`: the 5th argument `[no]` on every `PHP_ARG_WITH`/`PHP_ARG_ENABLE` call prevented `ext_shared` from being set, so `PHP_NEW_EXTENSION` registered the sources but never added the module to `PHP_MODULES`. `make` exited 0 having built nothing. This was the red "Extension loads" CI failure. - `config.m4`: `PHP_ADD_MAKEFILE_FRAGMENT` was called before `PHP_NEW_EXTENSION`; the fragment requires `ext_srcdir` (set by the latter), so Makefile.frag was silently never included. - `Makefile.frag`: NVCC objects are now prerequisites of the module target itself; previously a parallel make raced the link against the NVCC compilations. - `cuda.c`/`nvrtc.c`: real CUDA headers require `struct cudaDeviceProp` in C mode; `cudaMallocManaged` requires the flags argument; profiler functions need `cuda_profiler_api.h`. - `streams.c`: `cudaGraphInstantiate` signature differs between CUDA 11 (5 args) and CUDA 12 (3 args); now version-gated. - `nvrtc.c`: the driver API (libcuda) is no longer linked. It belongs to the NVIDIA driver, not the toolkit, and is absent on GPU-less build hosts and CI containers; it is now resolved lazily with `dlopen` at first use. The extension loads everywhere; only kernel compile/launch needs the driver. - `tensor.c`/`cuda.c`: `INIT_NS_CLASS_ENTRY` with an empty namespace registered the classes with a literal leading backslash; switched to `INIT_CLASS_ENTRY`. `CudaTensor::__toString` now declares its string return type (PHP 8.1+ warning). - Added `.gitignore` for phpize/configure/make artifacts. Verified: configure + parallel make + `php -d extension=cuda.so -m` succeed in GPU-less containers on CUDA 12.6 (PHP 8.1 and 8.4) and CUDA 11.8 (PHP 8.1); class and function registration confirmed. GPU execution of the .phpt suite remains the job of the self-hosted runner.
l4nos
force-pushed
the
feat/v0.2.0-driver-rebuild
branch
from
September 5, 2026 19:53
81abe41 to
292827a
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
This is a ground-up rebuild of the PHP CUDA extension into the driver layer that makes a "PHPTorch" technically possible: device-resident tensors with a real memory model, the NVIDIA library stack bound with correct semantics, streams/events/CUDA graphs, and NVRTC runtime kernel compilation, on a build system that actually compiles on modern toolchains.
It also fixes every root cause of issue #1 (Ubuntu 24.04 + CUDA 12 + cuDNN 9 build failure).
50 files changed.
Fixes #1
What was broken (found in a full code review)
The tree did not compile at all, on any platform:
cuda.creferenced module globals (allow_async_operations,default_device) that did not exist in the globals struct: hard compile errorcuda_matrix_multiply_kernel_wrapper()was declared extern, defined nowherePHP_FUNCTIONdeclarations inphp_cuda.hhad no implementations at all, while the test suite called several of themcompile.shdeletedconfig.m4via aconfig.*glob on every run, so the second build destroyed the project's build configurationconfig.m4had no working rule to compile.cufiles (Makefile.fragwas never included and referenced undefined variables)memory_pool.cuhusedstd::atomicwithout#include <atomic>Issue #1 root causes specifically:
-arch=sm_30: Kepler support was removed in CUDA 12; nvcc 12.x rejects itcudnnGetConvolutionForwardAlgorithm: removed in cuDNN 9nvtx3/); the build required-lnvToolsExtRuntime correctness bugs:
cuda_free()double-freed (explicitcudaFree+ resource destructor)cuda_matrix_multiply()had wrong dimension math (only worked for square matrices), no inner-dimension validation, and read an uninitialized variableneural_net.cuhad an out-of-bounds read and null derefscudaDeviceReset()ran unconditionally at module shutdown, which is hostile under php-fpm and persistent workersWhat this commit delivers
Build system (fixes #1)
config.m4rewritten:.cucompilation viaPHP_ADD_MAKEFILE_FRAGMENTwith working NVCC rules; GPU architectures auto-detected fromnvidia-smiwith toolkit-gated fallback lists (sm_70 through sm_120);--with-cuda-arch=override; cuDNN 7/8/9 detection withCUDNN_MAJORgating; NVTX3 (header-only) vs legacy nvToolsExt detection; NVRTC + driver API (-lnvrtc -lcuda) with stubs path handling; OpenMP optionalMakefile.fragrewritten: NVCC pattern rules, module-target prerequisitescompile.shrewritten: never touchesconfig.m4; refuses mixed toolchains (nvcc vsversion.json); warns when the driver-supported runtime is older than the toolkit;--uninstall;--test; PHP version check (8.1+)Dockerfile: Ubuntu 24.04 + CUDA 12.6 + cuDNN 9 reference environment.phptsuite; header self-containment checkCore runtime
E_WARNING+falseby default;cuda.error_mode=exceptionthrowsCudaExceptioncuda.default_device,cuda.error_mode,cuda.enable_cpu_fallback,cuda.enable_memory_poolCudaTensor substrate
CudaTensorPHP class: refcounted storage separate from views.reshape,transposeandsliceare zero-copy views sharing one allocation (the ATen/c10 memory model)add/sub/mul/div+ in-place variants),matmulvia cuBLAS with the row-major transpose trick (numerically correct, verified against reference values in tests)cuda_kernel_compile()/cuda_kernel_launch(): custom CUDA kernels authored in PHP userland, compiled to PTX at runtime, launched with CudaTensor/int/float arguments on any streamExecution model
cuda_graph_begin_capture/end_capture(instantiate) /launch/destroy: capture once, replay with near-zero CPU overheadNVIDIA library stack (partial)
cuda_cublas_matrix_multiply,cuda_cublas_gemm(alpha/beta),cuda_batch_gemm, all row-major correctcuda_cudnn_convolution_forwardwith_v7algorithm discovery (the legacy API was removed in cuDNN 9), workspace managementDead code removed
neural_net.*(superseded by the CudaTensor substrate),matrix_ops.*,tensor_ops.*,cuda_kernel.cu,cudnn_advanced.cuh,logger.cuh: all contained unimplemented or incorrect code referenced by nothing.Tests, docs, packaging
.phptfiles), including numeric verification of matmul/GEMM/convolution/softmax/broadcasting, an NVRTC kernel test, and a CUDA graph capture/replay test; tests skip gracefully without a GPU or optional componentsdocs/INSTALL-UBUNTU.md: exact Ubuntu 24.04 package set (including thenvidia-cuda-toolkitconflict warning from unable to build on Ubuntu 24.0 with cuda 12 etc #1) plus a troubleshooting tablecomposer.jsonwithphp-extconfigure options for PIEVerification
phpizepassesgcc -fsyntax-onlyagainst PHP 8.3 headersvoid**cast, missing<stdexcept>, 1-argcudaEventRecord, missingcuda_fp16.hinclude)./compile.sh --teston a CUDA host or the provided Docker image (docker run --gpus all)Known limitations (follow-up work)
CudaTensormatmul is 2-D only (batched matmul planned)fromBufferplanned